/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#include <cuqu/cuqu.h>
#ifdef CUPRINTF_INCLUDED
#include <utils/cuPrintf.cu>
#endif

typedef unsigned int uint;

#define BLOCK_SIZE 64
#define N_BLOCKS 1*14
#define NTHREADS (BLOCK_SIZE*N_BLOCKS)
#define WARP_SIZE 32

//#define EVENT_SIZE_W32 8 // <= warpSize !!!
//#define EVENT_SIZE_W32 WARP_SIZE // <= warpSize !!!

//#define QUEUE_EVENT_MAGIC 0xa5c4a5c4U
// 256 bits
typedef struct test_event {
    uint32_t w[8];
} test_event_t;

#define EVENT_MAGIC 0x10000

#define INIT_TEST_EVENT { 0,0,0,0,0,0,0,0 }

typedef cuqu::host::queue<test_event_t> queue_t;
typedef queue_t::event_t event_t;
typedef cuqu::cuda_queue_t cuda_queue_t;

//----------------------------------------------------------------------

//__device__ gpu_event_t event = INIT_GPU_EVENT(INIT_TEST_EVENT);
__device__ event_t event;

__global__ void dequeue(cuda_queue_t cuda_q, int num_events) 
{
    const unsigned int tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    const unsigned int bid = blockIdx.x + gridDim.x*blockIdx.y;
    const unsigned int block_size = blockDim.x*blockDim.y*blockDim.z;
    const unsigned int grid_size = gridDim.y*gridDim.x;
    //const unsigned int gid = tid + bid*block_size;

    int n = 0;
    int bad = 0;
    int wblk = 0;
    const int max_iter = 10000;
    int iter = max_iter;
    int retcode = 0;
    //event_t e;

#ifdef CUPRINTF_CU
    //cuPrintfRestrict(0, CUPRINTF_UNRESTRICTED);
    //cuPrintfRestrict(CUPRINTF_UNRESTRICTED, 0);
    //cuPrintf("num_events=%d\n", num_events);
#endif

    while(n < num_events) {
        retcode = cuda_q->timed_fetch(&event, 1000/*ms*/);
        if(cuqu::SUCCESS == retcode) {
            if(bid == 0 && tid == 0)
                cuqu_trace("event={%08x,%08x,%08x,...}\n", event->u.t.w[0], event->u.t.w[1], event->u.t.w[2]);

            // consistency check
#if defined(_DEBUG)
            if(n != event.u.t.w[0] || EVENT_MAGIC != event.u.t.w[1]) {
                if(tid % warpSize == 0) cuqu_error("[%d:%d] ERROR: n=%d != event={%d,0x%08x,0x%08x,...}\n", bid, tid, n, event.u.t.w[0], event.u.t.w[1], event.u.t.w[2]);
                break;
            }
            else
                cuqu_error("[%d:%d]  GOOD: n=%d    event={%d,0x%08x,0x%08x,...}\n", bid, tid, n, event.u.t.w[0], event.u.t.w[1], event.u.t.w[2]);
#endif
            iter = max_iter;
            ++n;
        } else if(retcode == cuqu::WOULDBLOCK) {
            ++wblk;
            cuqu_trace("ERROR: timeout exceeded for event %d\n", n);
            break;
        } else {
            ++bad;
            cuqu_error("[%d,%d] bad retcode=%d\n", bid, tid, retcode);
            break;
        }
    }
    //cuqu_trace("n=%d bad=%d wouldblock=%d retcode=%d event={%d, 0x%08x, ...}\n", n, bad, wblk, retcode, e->u.t.w[0], e->u.t.w[1]);
}

//----------------------------------------------------------------------

void push_n_events(queue_t *q, int n_pump_events)
{
    static int cnt = 0;
    for(int k=0; k<n_pump_events; ++k) {
        test_event_t e = {{cnt,EVENT_MAGIC,k,0x20202+k,0,0,0,0}};
        int ret = q->timed_push(&e, 1000);
        if(ret == cuqu::SUCCESS) {
            ++cnt;
        } else {
            printf("aborting at %d-th push\n", k); fflush(stdout);
            break;
        }
    }
}

//----------------------------------------------------------------------
void print_histogram(queue_t *q) {

#if CUQU_ENABLE_PERF_HIST
    unsigned int dev_hist[q->hist_size()];
    unsigned int *host_hist = q->get_host_hist_ptr();

    // Retrieve device histogram from device global memory.
    cuquSafeCall(cudaMemcpy(dev_hist, q->get_device_hist_ptr(), sizeof(unsigned int) * q->hist_size(), cudaMemcpyDeviceToHost));

    // print histogram
    for (size_t i = 0; i < q->hist_size(); i++) {
        printf("d[%2d]: %2d, h[%2d]: %2d\n", i, dev_hist[i], i, host_hist[i]);
    }
#else
    printf("performance histogram feature is disabled\n"); fflush(stdout);
#endif

}


//----------------------------------------------------------------------

void test(int nqevents, int n_pump_events)
{
#ifdef CUPRINTF_CU
#pragma message "calling cudaPrintfInit()"
    cudaPrintfInit();
#endif

#ifdef CUQU_DEBUG
    if(n_pump_events > nqevents) {
        n_pump_events = nqevents;
        printf("WARNING: debug build! forcing n_pump_events == nqevents\n"); fflush(stdout);
    }
#endif

    printf("sizeof(test_event_t)=%d\n", sizeof(test_event_t)); fflush(stdout);
    printf("barrier implementation: %s\n", CUQU_BARRIER_STR);

    printf("creating a queue with %d slots, pushing %d events\n", nqevents, n_pump_events); fflush(stdout);
    queue_t queue(nqevents, BLOCK_SIZE);
    queue_t *q = &queue;
    cuda_queue_t cuda_q = queue.get_cuda_ref();

    printf("pre-pushing %d events\n", nqevents);
    // WARN: correct only if pre-pushing events == queue size
    push_n_events(q, nqevents);

    cuquCheckMsg("aieee");

    // spawn GPU kernel
    printf("spawning dequeue() kernel, BLOCK_SIZE=%d N_BLOCKS=%d n_pump_events=%d\n", BLOCK_SIZE, N_BLOCKS, n_pump_events); fflush(stdout);
    dequeue<<<N_BLOCKS, BLOCK_SIZE>>>(cuda_q, n_pump_events);
        
    cuquCheckMsg("aieee");

    // push left out events
    int n_last_events = n_pump_events - nqevents;
    if(n_last_events > 0) {
        printf("pushing last %d event(s)\n", n_last_events); fflush(stdout);
        push_n_events(q, n_last_events);
    }
    cuquCheckMsg("aieee");

    printf("calling cudaThreadSynchronize()\n");
    cuquSafeCall(cudaThreadSynchronize());
    cuquCheckMsg("dequeue() execution failed");

    printf("kernel ended, reading back results\n"); fflush(stdout);

    cuqu::detail::raw_queue_t *raw_q = q->get_raw_ptr();
    printf("raw_queue: nr=%d nw=%d\n", raw_q->nr, raw_q->nw); fflush(stdout);

    print_histogram(q);

#ifdef CUPRINTF_CU
    cudaPrintfDisplay(stdout, true);
    cudaPrintfEnd();
#endif

    
    printf("destroying queue object\n"); fflush(stdout);
}

//----------------------------------------------------------------------

int main(int argc, char* argv[])
{
    cudaDeviceProp deviceProp;
    int idev;

    //
    int nqevents = 16;
    int n_pump_events = nqevents;

    // arg scan
    int opt;
    while ((opt = getopt(argc, argv, "hq:e:")) != -1) {
        switch (opt) {
        case 'q':
            nqevents = atoi(optarg);
            break;
        case 'e':
            n_pump_events = atoi(optarg);
            printf("parsed param -e %d\n", n_pump_events);
            break;
        default: /* '?' */
            fprintf(stderr, "Usage: %s [-e num_events][-q queue_size]\n",
                    argv[0]);
            exit(EXIT_FAILURE);
        }
    }

    // CUDA initialization
    printf("CUDA init\n"); fflush(stdout);
    idev = 0;
    cuquSafeCall(cudaSetDevice(idev));
    cuquSafeCall(cudaGetDeviceProperties(&deviceProp, idev));
    if(!deviceProp.canMapHostMemory) {
        fprintf(stderr, "Device %d cannot map host memory!\n", idev);
        exit(EXIT_FAILURE);
    }
    cuquSafeCall(cudaSetDeviceFlags(cudaDeviceMapHost));

#ifdef CUQU_DEBUG
    printf("WARNING: debug builds are not reliable, they may deadlock!\n"); fflush(stdout);
#endif

    test(nqevents, n_pump_events);

    //cudaThreadExit();    
  
    return EXIT_SUCCESS;
}

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
